#pragma OPENCL EXTENSION cl_arm_printf : enable
#pragma OPENCL EXTENSION cl_arm_integer_dot_product_int8 : enable

__constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;

#define TS 4

__kernel void matmul_int8_lbriobtp(__global char * out_buf, __global char4 *lhs_buf,
                                   __read_only image2d_t rhs_tex, int4 hwkt) {
    int height = hwkt.x, width = hwkt.y, kslice = hwkt.z/TS;
    int h = get_global_id(0) * TS;
    int w = get_global_id(1) * TS;

    // 203 ms
    int acc[TS][TS] = {0};
    for(int i=0; i<kslice; i+=TS) {
        char4 lhs = lhs_buf[h*kslice + i];
        acc[0] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+0))));
        acc[1] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+1))));
        acc[2] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+2))));
        acc[3] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+3))));

        // acc[4] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+4))));
        // acc[5] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+5))));
        // acc[6] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+6))));
        // acc[7] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+7))));

        // acc[8] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+8))));
        // acc[9] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+9))));
        // acc[10] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+10))));
        // acc[11] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+11))));
        
        // acc[12] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+12))));
        // acc[13] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+13))));
        // acc[14] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+14))));
        // acc[15] += arm_dot(lhs, convert_char4(read_imagei(rhs_tex, smp_none, (int2)(i, w+15))));
    }
    for (int j=0; j<4; j++)
        out_buf[h*width + w + j] = convert_char(acc[j] % 53);

    // float acc[32] = {0};
    // for(int i=0; i<kslice; i+=4) {
    //     float4 lhs = convert_float4(lhs_buf[h*kslice + i]);
    //     for (int j=0; j<16; j++) {
    //         acc[j] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+j))));
    //     }
    //     // acc[0] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+0))));
    //     // acc[1] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+1))));
    //     // acc[2] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+2))));
    //     // acc[3] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+3))));

    //     // acc[4] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+4))));
    //     // acc[5] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+5))));
    //     // acc[6] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+6))));
    //     // acc[7] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+7))));

    //     // acc[8] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+8))));
    //     // acc[9] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+9))));
    //     // acc[10] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+10))));
    //     // acc[11] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+11))));

    //     // acc[12] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+12))));
    //     // acc[13] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+13))));
    //     // acc[14] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+14))));
    //     // acc[15] += dot(lhs, convert_float4(read_imagei(rhs_tex, smp_none, (int2)(i, w+15))));
    // }
    // for (int j=0; j<16; j++)
    //     out_buf[h*width + w + j] = convert_char(((int)(acc[j])) % 53);
}